__kernel void scan(__global int * g_in,
				   __global int * g_out,
				   __global int * sums,
				   __local int *l_temp,
				   int n,
				   int block_number) {

	int thid = get_local_id(0);  
	int offset = 1;  

	l_temp[2*thid] = g_in[2*thid]; // load input into shared memory  
	l_temp[2*thid+1] = g_in[2*thid+1];  
 	
	for (int d = n>>1; d > 0; d >>= 1) // build sum in place up the tree  
	{   
		barrier(CLK_LOCAL_MEM_FENCE); 
		if (thid < d)  
		{  
			int i = offset*(2*thid+1)-1;  
			int j = offset*(2*thid+2)-1;  
			l_temp[j] += l_temp[i];  
		}  
		offset *= 2;
	}	
    
    
    
	// clear the last element  
	if (thid == 0) 
	{ 
		// for large array extension
    	sums[block_number] = l_temp[n-1];
		l_temp[n - 1] = 0; 
	} 
                 
	for (int d = 1; d < n; d *= 2) // traverse down tree & build scan  
	{  
		offset >>= 1;  
		barrier(CLK_LOCAL_MEM_FENCE);  
		if (thid < d)                       
		{   
			int i = offset*(2*thid+1)-1;  
			int j = offset*(2*thid+2)-1;  
 	
			float t = l_temp[j];  
			l_temp[j] += l_temp[i];  
			l_temp[i] = t;   
		}  
	}  
	barrier(CLK_LOCAL_MEM_FENCE);  

	g_out[2*thid] = l_temp[2*thid]; // write results to device memory  
	g_out[2*thid+1] = l_temp[2*thid+1];  
}


__kernel void scanSmall(__global int * g_in,
				   __global int * g_out,
				   __local int *l_temp,
				   int n) {

	int thid = get_local_id(0);  
	int offset = 1;  

	l_temp[2*thid] = g_in[2*thid]; // load input into shared memory  
	l_temp[2*thid+1] = g_in[2*thid+1];  
 	
	for (int d = n>>1; d > 0; d >>= 1) // build sum in place up the tree  
	{   
		barrier(CLK_LOCAL_MEM_FENCE); 
		if (thid < d)  
		{  
			int i = offset*(2*thid+1)-1;  
			int j = offset*(2*thid+2)-1;  
			l_temp[j] += l_temp[i];  
		}  
		offset *= 2;
	}	
    
    
    
	// clear the last element  
	if (thid == 0) 
	{ 
		l_temp[n - 1] = 0; 
	} 
                 
	for (int d = 1; d < n; d *= 2) // traverse down tree & build scan  
	{  
		offset >>= 1;  
		barrier(CLK_LOCAL_MEM_FENCE);  
		if (thid < d)                       
		{   
			int i = offset*(2*thid+1)-1;  
			int j = offset*(2*thid+2)-1;  
 	
			float t = l_temp[j];  
			l_temp[j] += l_temp[i];  
			l_temp[i] = t;   
		}  
	}  
	barrier(CLK_LOCAL_MEM_FENCE);  

	g_out[2*thid] = l_temp[2*thid]; // write results to device memory  
	g_out[2*thid+1] = l_temp[2*thid+1];  
}

__kernel void addBlockSums(__global int * scanedBlock,
				   __global int * blockSums,
				   __global int * output,
				   int blockId, 
				   int blockLength) {
	

	if(blockId > 0) 
	{
		for(int d = 0; d < blockLength; d+=1) 
		{
			output[d] = scanedBlock[d] + blockSums[blockId];
		} 	
	} else 
	{
		for(int d = 0; d < blockLength; d+=1) 
		{
			output[d] = scanedBlock[d];
		} 	
	}
}
